#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0
#define N 300000

__global__ void kernel_1()
{
    double sum = 0.0;
    for (int i = 0; i < N; i++)
        sum = sum + tan(0.1) * tan(0.1);
}
__global__ void kernel_2()
{
    double sum = 0.0;
    for (int i = 0; i < N; i++)
        sum = sum + tan(0.1) * tan(0.1);
}
__global__ void kernel_3()
{
    double sum = 0.0;
    for (int i = 0; i < N; i++)
        sum = sum + tan(0.1) * tan(0.1);
}
__global__ void kernel_4()
{
    double sum = 0.0;
    for (int i = 0; i < N; i++)
        sum = sum + tan(0.1) * tan(0.1);
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);

    // 建立一个长度为5的流数组
    int n_stream = 5;
    cudaStream_t *stream = (cudaStream_t *)malloc(n_stream * sizeof(cudaStream_t));
    for (int i = 0; i < n_stream; i++)
    {
        cudaStreamCreate(&stream[i]);
    }

    dim3 block(1);
    dim3 grid(1);

    cudaEvent_t start, stop;                                                    // 声明计时版本的event
    cudaEvent_t *event = (cudaEvent_t *)malloc(n_stream * sizeof(cudaEvent_t)); // 声明只用来同步的event数组
    for (int i = 0; i < n_stream; i++)
    {
        cudaEventCreateWithFlags(&event[i], cudaEventDisableTiming); // 只用来同步的event创建，去掉了计时功能
    }
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);         // 将start事件插入空流
    for (int i = 0; i < n_stream; i++) // 创建5个并发流
    {
        // 每个流上执行四个kernel
        kernel_1<<<grid, block, 0, stream[i]>>>();
        kernel_2<<<grid, block, 0, stream[i]>>>();
        kernel_3<<<grid, block, 0, stream[i]>>>();
        kernel_4<<<grid, block, 0, stream[i]>>>();
        cudaEventRecord(event[i], stream[i]); // 在每个流完成时插入不同的事件用来标记

        // 对最后一个流来说，由于event是插入到其他流操作的最后位置的
        // 这样只有等其他流工作都完成后，最后这个流才会开始执行
        cudaStreamWaitEvent(stream[n_stream - 1], event[i], 0);
    }
    cudaEventRecord(stop, 0); // 将stop事件插入空流
    CHECK(cudaEventSynchronize(stop));
    float elapsed_time;
    cudaEventElapsedTime(&elapsed_time, start, stop);
    std::cout << "space stream run time:" << elapsed_time << "ms" << std::endl;

    for (int i = 0; i < n_stream; i++)
    {
        cudaStreamDestroy(stream[i]);
    }
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    for (int i = 0; i < n_stream; i++)
    {
        cudaEventDestroy(event[i]);
    }
    free(stream);
    free(event);
    return 0;
}